#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 1;
class KernelLpNormV2
{
public:
  __aicore__ inline KernelLpNormV2() {}
  __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t typeKey,
                              uint32_t pType, float pValue, uint32_t totalLength, uint32_t stepSize, uint32_t unitCount)
  {
    ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
    this->typeKey = typeKey;
    this->pType = pType;
    this->pValue = pValue;
    if(typeKey == 0){
      this->yLength = 8;
    }else{
      this->yLength = 16;
    }
    this->totalLength = totalLength;
    this->stepSize = stepSize;
    this->unitCount = unitCount;

    this->blockLength = blockLength;
    this->tileNum =
        tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
    this->tileLength = tileLength / BUFFER_NUM;
    this->lasttileLength = lasttileLength;
    this->x_addr = x;
    xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x,
                        blockLength > this->totalLength * this->stepSize ? blockLength : this->totalLength * this->stepSize);
    yGm.SetGlobalBuffer((__gm__ DTYPE_X *)y,
                        blockLength > this->totalLength * this->stepSize ? blockLength : this->totalLength * this->stepSize);
    yGm_f.SetGlobalBuffer((__gm__ float *)y,
                        blockLength > this->totalLength * this->stepSize ? blockLength : this->totalLength * this->stepSize);
    if(stepSize == 1 && typeKey == 0 && pType == 0){
      this->tileLength = tileLength / 2;
      if(this->tileNum * 2 % 8 == 0){
        this->tileNumAlign = this->tileNum * 2;
      }else{
        this->tileNumAlign = (this->tileNum * 2 + 7) / 8 * 8;
      }
      if(this->lasttileLength % 16 == 0){
        this->lasttileLengthAlign = this->lasttileLength;
      }else{
        this->lasttileLengthAlign = (this->lasttileLength + 15) / 16 * 16;
      }
      pipe.InitBuffer(inQueueX_2, 2, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueT_2, 2, this->tileLength * sizeof(float));
      pipe.InitBuffer(outQueueY_2, 2, this->yLength * sizeof(float));
      pipe.InitBuffer(calcBuf, this->tileLength * sizeof(float));
    }else{
      pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
      pipe.InitBuffer(inQueueT, BUFFER_NUM, this->yLength * sizeof(float));
      pipe.InitBuffer(outQueueY, BUFFER_NUM, this->yLength * sizeof(DTYPE_X));
      pipe.InitBuffer(outQueueY_F, BUFFER_NUM, this->yLength * sizeof(float));
      pipe.InitBuffer(calcBuf, BUFFER_NUM * this->tileLength * sizeof(float));
      pipe.InitBuffer(xfBuf, BUFFER_NUM * this->tileLength * sizeof(float));
      pipe.InitBuffer(calcBuf1, BUFFER_NUM * this->tileLength * sizeof(uint8_t));
      pipe.InitBuffer(calcBuf2, BUFFER_NUM * this->tileLength * sizeof(float));
    }
  }
  __aicore__ inline void Process()
  {
    if(stepSize != 1){
      Process3();
      return ;
    }
    if(typeKey == 0 && pType == 0){
      Process2();
      return ;
    }
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++)
    {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
    }
    if(this->pType == 0){
      // ========CopyIn======
      LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
      DataCopy(tLocal, yGm_f[0], this->yLength);
      inQueueT.EnQue(tLocal);
      if(this->typeKey == 1){
        // ========Compute======
        LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
        tLocal = inQueueT.DeQue<float>();
        Sqrt(tLocal, tLocal, this->yLength);
        Cast(yLocal, tLocal, RoundMode::CAST_NONE, this->yLength);
        // Duplicate(yLocal, (DTYPE_X)1.2, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<DTYPE_X>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<DTYPE_X>();
        DataCopy(yGm[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }else{
        // ========Compute======
        LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
        tLocal = inQueueT.DeQue<float>();
        Sqrt(yLocal, tLocal, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<float>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<float>();
        DataCopy(yGm_f[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }
    }else if(this->pType == 1 || this->pType == 2 || this->pType == 3 || this->pType == 4){
      if(this->typeKey == 1){
        // ========CopyIn======
        LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
        DataCopy(tLocal, yGm_f[0], this->yLength);
        inQueueT.EnQue(tLocal);
        // ========Compute======
        tLocal = inQueueT.DeQue<float>();
        LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
        Cast(yLocal, tLocal, RoundMode::CAST_NONE, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<DTYPE_X>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<DTYPE_X>();
        DataCopy(yGm[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }else{
        // ========CopyIn======
        LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
        DataCopy(tLocal, yGm_f[0], this->yLength);
        inQueueT.EnQue(tLocal);
        // ========Compute======
        tLocal = inQueueT.DeQue<float>();
        LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
        Adds(yLocal, tLocal, (float)0.0, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<float>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<float>();
        DataCopy(yGm_f[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }
      
    }else {
      if(this->typeKey == 1){
        // ========CopyIn======
        LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
        DataCopy(tLocal, yGm_f[0], this->yLength);
        inQueueT.EnQue(tLocal);
        // ========Compute======
        LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
        tLocal = inQueueT.DeQue<float>();
        Ln(tLocal, tLocal, this->tileLength);
        Muls(tLocal, tLocal, (float)(1.0f/this->pValue), this->tileLength);
        Exp(tLocal, tLocal, this->tileLength);
        Cast(yLocal, tLocal, RoundMode::CAST_NONE, this->yLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<DTYPE_X>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<DTYPE_X>();
        DataCopy(yGm[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }else{
        // ========CopyIn======
        LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
        DataCopy(tLocal, yGm_f[0], this->yLength);
        inQueueT.EnQue(tLocal);
        // ========Compute======
        LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
        tLocal = inQueueT.DeQue<float>();
        Ln(tLocal, tLocal, this->tileLength);
        Muls(tLocal, tLocal, (float)(1.0f/this->pValue), this->tileLength);
        Exp(yLocal, tLocal, this->tileLength);
        inQueueT.FreeTensor(tLocal);
        outQueueY.EnQue<float>(yLocal);
        // ========CopyOut======
        yLocal = outQueueY.DeQue<float>();
        DataCopy(yGm_f[0], yLocal, this->yLength);
        outQueueY.FreeTensor(yLocal);
      }
    }
  }

  __aicore__ inline void Process3(){
    int32_t loopCount = this->tileNum;
    for(int32_t k = 0; k < this->unitCount; k++){
      for(int32_t j = 0; j < this->stepSize; j++){
        for (int32_t i = 0; i < loopCount; i++)
        {
          CopyIn3(i, j, k);
          Compute(i);
          CopyOut3(i, j, k);
        }
        if(this->pType == 0){
          // ========CopyIn======
          LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
          DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
          inQueueT.EnQue(tLocal);
          if(this->typeKey == 1){
            // ========Compute======
            LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
            tLocal = inQueueT.DeQue<float>();
            Sqrt(tLocal, tLocal, this->yLength);
            Cast(yLocal, tLocal, RoundMode::CAST_NONE, this->yLength);
            // Duplicate(yLocal, (DTYPE_X)1.2, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<DTYPE_X>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<DTYPE_X>();
            DataCopy(yGm[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }else{
            // ========Compute======
            LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
            tLocal = inQueueT.DeQue<float>();
            Sqrt(yLocal, tLocal, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<float>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<float>();
            DataCopy(yGm_f[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }
        }else if(this->pType == 1 || this->pType == 2 || this->pType == 3 || this->pType == 4){
          if(this->typeKey == 1){
            // ========CopyIn======
            LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
            DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
            inQueueT.EnQue(tLocal);
            // ========Compute======
            tLocal = inQueueT.DeQue<float>();
            LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
            Cast(yLocal, tLocal, RoundMode::CAST_NONE, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<DTYPE_X>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<DTYPE_X>();
            DataCopy(yGm[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }else{
            // ========CopyIn======
            LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
            DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
            inQueueT.EnQue(tLocal);
            // ========Compute======
            tLocal = inQueueT.DeQue<float>();
            LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
            Adds(yLocal, tLocal, (float)0.0, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<float>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<float>();
            DataCopy(yGm_f[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }
          
        }else {
          if(this->typeKey == 1){
            // ========CopyIn======
            LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
            DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
            inQueueT.EnQue(tLocal);
            // ========Compute======
            LocalTensor<DTYPE_X> yLocal = outQueueY.AllocTensor<DTYPE_X>();
            tLocal = inQueueT.DeQue<float>();
            Ln(tLocal, tLocal, this->tileLength);
            Muls(tLocal, tLocal, (float)(1.0f/this->pValue), this->tileLength);
            Exp(tLocal, tLocal, this->tileLength);
            Cast(yLocal, tLocal, RoundMode::CAST_NONE, this->yLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<DTYPE_X>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<DTYPE_X>();
            DataCopy(yGm[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }else{
            // ========CopyIn======
            LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
            DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
            inQueueT.EnQue(tLocal);
            // ========Compute======
            LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
            tLocal = inQueueT.DeQue<float>();
            Ln(tLocal, tLocal, this->tileLength);
            Muls(tLocal, tLocal, (float)(1.0f/this->pValue), this->tileLength);
            Exp(yLocal, tLocal, this->tileLength);
            inQueueT.FreeTensor(tLocal);
            outQueueY.EnQue<float>(yLocal);
            // ========CopyOut======
            yLocal = outQueueY.DeQue<float>();
            DataCopy(yGm_f[j + k * this->stepSize], yLocal, this->yLength);
            outQueueY.FreeTensor(yLocal);
          }
        }
      }
    }
    
  }

  __aicore__ inline void Process2() // sqrt(sum(x^2))
  {
    int32_t loopCount = this->tileNum * 2;
    LocalTensor<float> tLocal = inQueueT_2.AllocTensor<float>();
    inQueueT_2.EnQue<float>(tLocal);
    for (int32_t progress = 0; progress < loopCount; progress++)
    {
      CopyIn2(progress);
      Compute2(progress);
      CopyOut2(progress);
    }
    // Compute
    tLocal = inQueueT_2.DeQue<float>();
    LocalTensor<float> yLocal = outQueueY_2.AllocTensor<float>();
    LocalTensor<float> wLocal = calcBuf.Get<float>();
    ReduceSum(yLocal, tLocal, wLocal, this->tileLength);
    Sqrt(yLocal, yLocal, this->yLength);
    inQueueT_2.FreeTensor(tLocal);
    outQueueY_2.EnQue(yLocal);
    // CopyOut
    yLocal = outQueueY_2.DeQue<float>();
    DataCopy(yGm_f[0], yLocal, this->yLength);
    outQueueY_2.FreeTensor(yLocal);
  }

private:
  __aicore__ inline void CopyIn2(int32_t progress)
  {
    LocalTensor<DTYPE_X> xLocal = inQueueX_2.AllocTensor<DTYPE_X>();
    if (progress == this->tileNum * 2 - 2) {
      DataCopy(
          xLocal[0],
          xGm[progress  * this->tileLength],
          this->lasttileLengthAlign / 2);
    } else if (progress == this->tileNum * 2 - 1) {
      DataCopy(xLocal[0], xGm[(progress - 1) * this->tileLength + this->lasttileLengthAlign / 2],
              this->lasttileLengthAlign / 2);
    } else{
      DataCopy(xLocal[0], xGm[progress * this->tileLength],
              this->tileLength);
    }
    inQueueX_2.EnQue(xLocal);
  }
  __aicore__ inline void Compute2(int32_t progress)
  {
    LocalTensor<float> tLocal = inQueueT_2.DeQue<float>();
    LocalTensor<float> xLocal_f = inQueueX_2.DeQue<float>();
    LocalTensor<float> yLocal = outQueueY_2.AllocTensor<float>();
    if(progress >= this->tileNum * 2 - 2){
      if(progress == 0){
        Mul(tLocal, xLocal_f, xLocal_f, this->lasttileLengthAlign / 2);
      }else{
        Mul(xLocal_f, xLocal_f, xLocal_f, this->lasttileLengthAlign / 2);
        Add(tLocal, tLocal, xLocal_f, this->lasttileLengthAlign / 2);
      }
    }else{
      if(progress == 0){
        Mul(tLocal, xLocal_f, xLocal_f, this->tileLength);
      }else{
        Mul(xLocal_f, xLocal_f, xLocal_f, this->tileLength);
        Add(tLocal, tLocal, xLocal_f, this->tileLength);
      }
    }
    inQueueT_2.EnQue<float>(tLocal);
    outQueueY_2.EnQue(yLocal);
    inQueueX_2.FreeTensor(xLocal_f);
  }
  __aicore__ inline void CopyOut2(int32_t progress)
  {
    LocalTensor<float> yLocal = outQueueY_2.DeQue<float>();
    outQueueY_2.FreeTensor(yLocal);
  }
  __aicore__ inline void CopyIn(int32_t progress)
  {
    LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
    LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
          //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
          DataCopy(xLocal[0], xGm[0], this->tileLength);
          for(int i = this->totalLength; i < this->tileLength; i++){
            if(pType == 0 || pType == 3 || pType == 4 || pType == 5){
              xLocal.SetValue(i, (DTYPE_X)0);
            }else{
              xLocal.SetValue(i, xLocal.GetValue(0));
            }
          }
        } else {
          //将最后一个分块的起始地址向前移动tileLength-lasttileLength
          DataCopy(
              xLocal[0],
              xGm[(progress - 1) * this->tileLength + this->lasttileLength],
              this->tileLength);
          for(int i = 0; i < this->tileLength - this->lasttileLength; i++){
            if(pType == 0 || pType == 3 || pType == 4 || pType == 5){
              xLocal.SetValue(i, (DTYPE_X)0);
            }else{
              xLocal.SetValue(i, xLocal.GetValue(this->tileLength - this->lasttileLength));
            }
          }
        }
      } else {
        DataCopy(xLocal[0], xGm[progress * this->tileLength],
                 this->tileLength);
        
      }
    }
    DataCopy(tLocal, yGm_f[0], this->yLength);
    if(progress == 0){
      tLocal.SetValue(0, (DTYPE_X)0);
    }
    inQueueT.EnQue(tLocal);
    inQueueX.EnQue(xLocal);
    
  }
  __aicore__ inline void Compute(int32_t progress)
  {
    LocalTensor<float> yLocal = outQueueY_F.AllocTensor<float>();
    LocalTensor<float> tLocal = inQueueT.DeQue<float>();
    LocalTensor<float> wLocal = calcBuf.Get<float>();
    
    if(this->typeKey == 0){
      LocalTensor<float> xLocal = inQueueX.DeQue<float>();
      LocalTensor<float> xLocal_f = xLocal;
      
      if(pType == 0){
        //sqrt(sum(abs(x)^2))   先求abs(x)^2
        Mul(xLocal_f, xLocal_f, xLocal_f, this->tileLength);
        ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 1){
        //max(abs(x))
        Abs(xLocal_f, xLocal_f, this->tileLength);
        ReduceMax(yLocal, xLocal_f, wLocal, this->tileLength, false);
        Max(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 2){
        //min(abs(x))
        if(progress < 1){
          Abs(tLocal, xLocal_f, this->yLength);
        }
        Abs(xLocal_f, xLocal_f, this->tileLength);
        ReduceMin(yLocal, xLocal_f, wLocal, this->tileLength, false);
        Min(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 3){
        // sum(x!=0) 即不等于0的个数
        LocalTensor<uint8_t> resLocal = calcBuf1.Get<uint8_t>();
        LocalTensor<float> zeroLocal = calcBuf2.Get<float>();
        
        Duplicate(zeroLocal, (float)0.0, this->tileLength);
        Compare(resLocal, xLocal_f, zeroLocal, CMPMODE::EQ, this->tileLength);
        Select(zeroLocal, resLocal, zeroLocal, static_cast<float>(1), SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);
        ReduceSum(yLocal, zeroLocal, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 4){
        // sum(abs(x))
        Abs(xLocal_f, xLocal_f, this->tileLength);
        ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 5){
        //sum(abs(x)^p)^(1/p)
        Abs(xLocal_f, xLocal_f, this->tileLength);
        Ln(xLocal_f, xLocal_f, this->tileLength);
        Muls(xLocal_f, xLocal_f, (float)(this->pValue), this->tileLength);
        Exp(xLocal_f, xLocal_f, this->tileLength);
        ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }
      outQueueY_F.EnQue<float>(yLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueT.FreeTensor(tLocal);
    }else{
      LocalTensor<half> xLocal = inQueueX.DeQue<half>();
      LocalTensor<float> xLocal_f = xfBuf.Get<float>();
      Cast(xLocal_f, xLocal, RoundMode::CAST_NONE, this->tileLength);
      if(pType == 0){
        //sqrt(sum(abs(x)^2))   先求abs(x)^2
        Abs(xLocal_f, xLocal_f, this->tileLength);
        Mul(xLocal_f, xLocal_f, xLocal_f, this->tileLength);
        ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 1){
        //max(abs(x))
        Abs(xLocal_f, xLocal_f, this->tileLength);
        ReduceMax(yLocal, xLocal_f, wLocal, this->tileLength, false);
        Max(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 2){
        //min(abs(x))
        if(progress < 1){
          Abs(tLocal, xLocal_f, this->yLength);
        }
        Abs(xLocal_f, xLocal_f, this->tileLength);
        ReduceMin(yLocal, xLocal_f, wLocal, this->tileLength, false);
        Min(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 3){
        // sum(x!=0) 即不等于0的个数
        LocalTensor<uint8_t> resLocal = calcBuf1.Get<uint8_t>();
        LocalTensor<float> zeroLocal = calcBuf2.Get<float>();
        
        Duplicate(zeroLocal, (float)0.0, this->tileLength);
        Compare(resLocal, xLocal_f, zeroLocal, CMPMODE::EQ, this->tileLength);
        Select(zeroLocal, resLocal, zeroLocal, static_cast<float>(1), SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);
        ReduceSum(yLocal, zeroLocal, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 4){
        // sum(abs(x))
        Abs(xLocal_f, xLocal_f, this->tileLength);
        ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }else if(pType == 5){
        //sum(abs(x)^p)^(1/p)
        Abs(xLocal_f, xLocal_f, this->tileLength);
        Ln(xLocal_f, xLocal_f, this->tileLength);
        Muls(xLocal_f, xLocal_f, (float)(this->pValue), this->tileLength);
        Exp(xLocal_f, xLocal_f, this->tileLength);
        ReduceSum(yLocal, xLocal_f, wLocal, this->tileLength);
        Add(yLocal, yLocal, tLocal, this->yLength);
      }
      outQueueY_F.EnQue<float>(yLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueT.FreeTensor(tLocal);
    }
    
  }
  __aicore__ inline void CopyOut(int32_t progress)
  {
    LocalTensor<float> yLocal = outQueueY_F.DeQue<float>();
    DataCopy(yGm_f[0], yLocal, this->yLength);
    outQueueY_F.FreeTensor(yLocal);
  }

  __aicore__ inline void CopyIn3(int32_t progress, int j, int k)
  {
    LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
    LocalTensor<float> tLocal = inQueueT.AllocTensor<float>();
    uint32_t startIndex = j + k * this->stepSize * this->totalLength;
    if (progress == this->tileNum - 1) {
      for(int i = 0; i < lasttileLength; i++){
        xLocal.SetValue(i, ((__gm__ DTYPE_X *)x_addr)[startIndex + progress * tileLength * this->stepSize + i * stepSize]);
      }
      for(int i = lasttileLength; i < this->tileLength; i++){
        if(pType == 0 || pType == 3 || pType == 4 || pType == 5){
          xLocal.SetValue(i, (DTYPE_X)0);
        }else{
          xLocal.SetValue(i, xLocal.GetValue(0));
        }
      }
    } else {
      for(int i = 0; i < tileLength; i++){
        xLocal.SetValue(i, ((__gm__ DTYPE_X *)x_addr)[startIndex + progress * tileLength * this->stepSize + i * stepSize]);
      }
    }
    DataCopy(tLocal, yGm_f[j + k * this->stepSize], this->yLength);
    if(progress == 0){
      tLocal.SetValue(0, (DTYPE_X)0);
    }
    inQueueT.EnQue(tLocal);
    inQueueX.EnQue(xLocal);
  }

  __aicore__ inline void CopyOut3(int32_t progress, int j, int k)
  {
    LocalTensor<float> yLocal = outQueueY_F.DeQue<float>();
    DataCopy(yGm_f[j + k * this->stepSize], yLocal, this->yLength);
    outQueueY_F.FreeTensor(yLocal);
  }

private:
  TPipe pipe;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueT;
  TQue<QuePosition::VECIN, 2> inQueueX_2, inQueueT_2;
  TQue<QuePosition::VECOUT, 2> outQueueY_2;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY, outQueueY_F;
  TBuf<TPosition::VECCALC> calcBuf, calcBuf1, calcBuf2, xfBuf;
  GlobalTensor<DTYPE_X> xGm;
  GlobalTensor<float> yGm_f;
  GlobalTensor<DTYPE_Y> yGm;
  GM_ADDR x_addr;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileNumAlign;
  uint32_t tileLength;
  uint32_t totalLength;
  uint32_t stepSize;
  uint32_t unitCount;
  uint32_t lasttileLength;
  uint32_t lasttileLengthAlign;
  uint32_t typeKey;
  uint32_t pType;
  uint32_t yLength;
  float pValue;
};
extern "C" __global__ __aicore__ void lp_norm_v2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    
    KernelLpNormV2 op;

    op.Init(x, y, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.typeKey,
            tiling_data.pType, tiling_data.pValue, tiling_data.totalLength, tiling_data.stepSize, tiling_data.unitCount);
    op.Process();
    // TODO: user kernel impl
}
#ifndef __CCE_KT_TEST__
void lp_norm_v2_do(uint32_t blockDim, void* l2ctrl, void* stream,
                       uint8_t* x, uint8_t* y,
                       uint8_t* workspace, uint8_t* tiling) {
  
  lp_norm_v2<<<blockDim, l2ctrl, stream>>>(x, y, workspace, tiling);
}
#endif